#ifndef CPU_ONLY
#ifdef USE_PROF

#include <algorithm>
#include <limits>

#include <cmath>
#include <cctype>
#include <ctime>

#include <cuda.h>
#include <cupti.h>

#include <boost/thread.hpp>
#include <boost/make_shared.hpp>
#include <boost/thread/mutex.hpp>

#include <glog/logging.h>

#include "caffe/util/async_tracker.hpp"
#include "caffe/util/info_log.hpp"

#define NS2MS 1000000.0

#define CHECK_KERNEL_RECORD(record) {                                   \
     if (record->start == 0 or record->end == 0) { \
       LOG(FATAL) << "WARNING! Cannot record the START timestamp of the " \
                  << record->name << " kernel. [start, end] = [" << record->start \
                  << ", " << record->end << "]."; \
     }                                  \
  }

#define CUPTI_CHECK(condition)                                          \
  do {                                                                  \
    CUptiResult status = condition;                                     \
    if (status != CUPTI_SUCCESS) {                                      \
      const char *errstr;                                               \
      cuptiGetResultString(status, &errstr);                            \
      LOG(FATAL) << __FILE__ << ":" << __LINE__ << ": error: "          \
                 << "function " << #condition << " failed with error "  \
                 << errstr << ".";                                      \
      exit(-1);                                                         \
    }                                                                   \
  } while (0)

#define CUPTI_CHECK_ARG(condition, arg1)                              \
  do {                                                                \
    CUptiResult status = condition;                                   \
    if (status != CUPTI_SUCCESS) {                                    \
      const char *errstr;                                             \
      cuptiGetResultString(status, &errstr);                          \
      LOG(FATAL) << __FILE__ << ":" << __LINE__ << ": error:  "       \
                 << "function " << #condition << " failed with error" \
                 << errstr << "(" << arg1 << ").";                    \
      exit(-1);                                                       \
    }                                                                 \
  } while (0)

namespace caffe {
  using std::string;
  using std::vector;
  using std::fstream;
  using std::time_t;
  using std::tm;
  using sstream = std::stringstream;

  static boost::mutex s_prof_mutex;
  /**
   * Member variable used to identify whether start profiling.
   * Used in the BufferCompleted callback function.
   */
  static bool s_profiler_flag = false;

  // Global static ActivityTracker object. Maybe can be designed with boost::mutex.
  ActivityTracker::ActivityTrackerProxy ActivityTracker::s_activity_tracker_proxy_;

  inline void str_tolower(string& str) {
    std::transform(str.begin(), str.end(), str.begin(),
                   [](unsigned char c) -> unsigned char { return std::tolower(c); }
                   );
  }

  template <typename T>
  bool Compare(T v1, T v2) {
    return (v1 < v2);
  }

  string GetCurrTime() {
    time_t curr_time = time(NULL);
    tm *local_time = localtime(&curr_time);
    sstream temp_ss;
    temp_ss << local_time->tm_year << "-" << local_time->tm_mon << "-" << local_time->tm_mday;
    string result = temp_ss.str();
    temp_ss.str("");
    temp_ss.clear();

    return result;
  }

  const char* getActivityOverheadKindString(CUpti_ActivityOverheadKind kind) {
    switch (kind) {
    case CUPTI_ACTIVITY_OVERHEAD_DRIVER_COMPILER: // Compiler(JIT) overhead
      return "COMPILER OVERHEAD";
    case CUPTI_ACTIVITY_OVERHEAD_CUPTI_BUFFER_FLUSH: // Activity buffer flush overhead
      return "BUFFER FLUSH OVERHEAD";
    case CUPTI_ACTIVITY_OVERHEAD_CUPTI_INSTRUMENTATION: // CUPTI instrumentation overhead
      return "INSTRUMENTATION OVERHEAD";
    case CUPTI_ACTIVITY_OVERHEAD_CUPTI_RESOURCE: // CUPTI resource creation and destruction overhead
      return "RESOURCE";
    default:
      break;
    }

    return "<UNKNOWN>";
  }

  const char* getActivityObjectKindString(CUpti_ActivityObjectKind kind) {
    switch (kind) {
    case CUPTI_ACTIVITY_OBJECT_PROCESS:
      return "PROCESS";
    case CUPTI_ACTIVITY_OBJECT_THREAD:
      return "THREAD";
    case CUPTI_ACTIVITY_OBJECT_DEVICE:
      return "DEVICE";
    case CUPTI_ACTIVITY_OBJECT_CONTEXT:
      return "CONTEXT";
    case CUPTI_ACTIVITY_OBJECT_STREAM:
      return "STREAM";
    default:
      break;
    }

    return "<UNKNOWN>";
  }

  /**
   * @brief      Kernel search function.
   *
   * @details    This function is used to check whether there is an identical
   *             kernel has been recorded.
   *
   * @param      kernel     Kernel record to be processed.
   *
   * @return     int   Index if the kernel is found in the kernel_vec_ptr_, and -1, otherwise.
   */
  int FindKernelConfig(boost::shared_ptr<vector<Kernel_t>> kernels_vec_ptr, const Kernel_t& kernel) {
    for (int i = 0; i < kernels_vec_ptr->size(); ++ i) {
      if (kernels_vec_ptr->at(i) == kernel) {
        return i;
      }
    }

    return -1;
  }

  /**
   * @brief      Buffer request function
   *
   * @details    Function used to handle buffer requests from CUPTI runtime. This function will allocate
   *             buffers for CUPTI to store kernel records.
   *
   * @param[out]  buffer      Buffer allocated for profiling.
   * @param[out]  size        Size of the buffer allocated.d
   * @param[out]  max_num_records  The maximum number of records that should be placed in the buffer. 0
   *                               represents that the buffer is filled with as many records as possible.
   */
  void CUPTIAPI BufferRequested(uint8_t **buffer, size_t *size, size_t *max_records) {
    // If s_profiler_flag is set to false, then buffer will be set to nullptr, which represents
    // no kernels will be recorded.
    if (s_profiler_flag) {
      uint8_t *temp_buf = new uint8_t[BUFFER_SIZE + ALIGN_SIZE];

      if (temp_buf == nullptr) {
        LOG(FATAL) << "Error! Cannot allocate new buffer for activity record collecting! OUT OF MEMORY!";
      }

      *size = BUFFER_SIZE;
      *buffer = ALIGN_BUFFER(temp_buf, ALIGN_SIZE);
      *max_records = 0;

      if (ActivityTracker::Get().GetBuffOverheadRecord()) {
        sstream temp_ss;
        temp_ss << (BUFFER_SIZE + ALIGN_SIZE) * sizeof(uint8_t) / 1024 << " KB" << std::endl;
        InfoLog::Get().RecordInfoLog("CuptiOverhead", "cupti_buffer", "CUPTI-BUFFER", temp_ss.str());
        temp_ss.str("");
        temp_ss.clear();
      }
    } else {
      *size = 0;
      *max_records = 0;
      *buffer = nullptr;
    }
  }

  /**
   * @brief      Buffer completion function.
   *
   * @details    Function used by CUPTI to return a buffer of activity records. The buffer contains valid_size
   *             bytes of activity krecords which should be read using cuptiActivityGetNextRecord.
   *
   * @param[in]  context    The context this buffer is associated with. If NULL, the buffer is associated
   *               with the global activities. This field is deprecated as of CUDA 6.0 and will always be NULL.
   * @param[in]  stream_id  CUDA stream ID.
   * @param[in]  buffer     The activity record buffer.
   * @param[in]  size
   * @param[in]  valid_size Size of activity records in bytes.
   */
  void CUPTIAPI BufferCompleted(CUcontext ctx, uint32_t stream_id, uint8_t *buffer, size_t size, size_t valid_size) {
    CUptiResult cupti_status;
    CUpti_Activity *record = nullptr;

    if (valid_size > 0) {
      do {
        // Get the next kernel recorded if exists.
        cupti_status = cuptiActivityGetNextRecord(buffer, valid_size, &record);
        if (cupti_status == CUPTI_SUCCESS) {
          // Parse the kernel execution information.
          ParseCUptiActivity(record, ActivityTracker::Get());
        } else if (cupti_status == CUPTI_ERROR_MAX_LIMIT_REACHED) {
          break;
        } else {
          CUPTI_CHECK(cupti_status);
        }
      } while (true);

      // LOG the number of records dropped if exists.
      size_t dropped_records;
      CUPTI_CHECK(cuptiActivityGetNumDroppedRecords(ctx, stream_id, &dropped_records));
      if (dropped_records != 0) {
        LOG(INFO) << "Dropped " << dropped_records << " activity records!";
      }
    }

    // Free the allocated buffer.
    if (buffer != nullptr) {
      delete[] buffer;
    }
  }

  void CUPTIAPI MetricValueCallback(void *userdata, CUpti_CallbackDomain domain, CUpti_CallbackId cbid, const CUpti_CallbackData *cbInfo) {
    Metric_t *metricData = reinterpret_cast<Metric_t*>(userdata);
    unsigned int i, j, k;

    // This callback is enabled only for launch so we shouldn't see anything else.
    if ((cbid != CUPTI_RUNTIME_TRACE_CBID_cudaLaunch_v3020) and
        (cbid != CUPTI_RUNTIME_TRACE_CBID_cudaLaunchKernel_v7000)) {
      LOG(FATAL) << "Unexpected cbid " << cbid;
    }

    // on entry, enable all the event groups being collected this pass.
    // for metrics we collect for all instances of the event.
    if (cbInfo->callbackSite == CUPTI_API_ENTER) {
      cudaDeviceSynchronize();

      CUPTI_CHECK(cuptiSetEventCollectionMode(cbInfo->context,
                                              CUPTI_EVENT_COLLECTION_MODE_KERNEL));

      for (i = 0; i < metricData->event_groups_->numEventGroups; ++ i) {
        uint32_t all = 1;
        CUPTI_CHECK(cuptiEventGroupSetAttribute(metricData->event_groups_->eventGroups[i],
                                                CUPTI_EVENT_GROUP_ATTR_PROFILE_ALL_DOMAIN_INSTANCES,
                                                sizeof(all), &all));
        CUPTI_CHECK(cuptiEventGroupEnable(metricData->event_groups_->eventGroups[i]));
      }
    }

    // on exit, read and record event values.
    if (cbInfo->callbackSite == CUPTI_API_EXIT) {
      cudaDeviceSynchronize();

      // for each group, read the event values from the group and record din metricData.
      for (i = 0; i < metricData->event_groups_->numEventGroups; ++ i) {
        CUpti_EventGroup group = metricData->event_groups_->eventGroups[i];
        CUpti_EventDomainID groupDomain;
        CUpti_EventID *eventIds;
        uint32_t numEvents, numInstances = 0, numTotalInstances;
        size_t groupDomainSize = sizeof(groupDomain);
        size_t numEventsSize = sizeof(numEvents);
        // size_t numIntancesSize = sizeof(numInstances);
        size_t numTotalInstancesSize = sizeof(numTotalInstances);
        uint64_t *values, normalized, *sum;
        size_t valueSize, eventIdsSize;
        size_t numCountersRead = 0;

        CUPTI_CHECK(cuptiEventGroupGetAttribute(group,
                                                CUPTI_EVENT_GROUP_ATTR_EVENT_DOMAIN_ID,
                                                &groupDomainSize, &groupDomain));
        CUPTI_CHECK(cuptiDeviceGetEventDomainAttribute(metricData->device_, groupDomain,
                                                       CUPTI_EVENT_DOMAIN_ATTR_TOTAL_INSTANCE_COUNT,
                                                       &numTotalInstancesSize, &numTotalInstances));
        CUPTI_CHECK(cuptiEventGroupGetAttribute(group,
                                                CUPTI_EVENT_GROUP_ATTR_INSTANCE_COUNT,
                                                &numEventsSize, &numEvents));
        eventIdsSize = numEvents * sizeof(CUpti_EventID);
        eventIds = static_cast<CUpti_EventID *>(malloc(eventIdsSize));
        CUPTI_CHECK(cuptiEventGroupGetAttribute(group,
                                                CUPTI_EVENT_GROUP_ATTR_EVENTS,
                                                &eventIdsSize, eventIds));

        valueSize = sizeof(uint64_t) * numInstances * numEvents;
        values = static_cast<uint64_t*>(malloc(valueSize));
        CUPTI_CHECK(cuptiEventGroupReadAllEvents(group,
                                                 CUPTI_EVENT_READ_FLAG_NONE,
                                                 &valueSize, values,
                                                 &eventIdsSize, eventIds,
                                                 &numCountersRead));

        if (metricData->event_idx_ >= metricData->num_events_) {
          LOG(FATAL) << "Error! Too many events collected, metric expects only "
                     << static_cast<int>(metricData->num_events_);
        }

        sum = static_cast<uint64_t*>(calloc(sizeof(uint64_t), numEvents));
        // sum collect event values from all instances.
        for (k = 0; k < numInstances; ++ k) {
          for (j = 0; j < numEvents; ++ j) {
            sum[j] += values[(k * numEvents) + j];
          }
        }

        for (j = 0; j < numEvents; ++ j) {
          // normalize the event value to represent the total number of domain instances
          // on the device.
          normalized = (sum[j] * numTotalInstances) / numInstances;

          metricData->event_id_array_[metricData->event_idx_] = eventIds[j];
          metricData->event_val_array_[metricData->event_idx_] = normalized;
        }

        free(values);
        free(sum);
      }

      for (i = 0; i < metricData->event_groups_->numEventGroups; ++ i) {
        CUPTI_CHECK(cuptiEventGroupDisable(metricData->event_groups_->eventGroups[i]));
      }
    }
  }

  void ParseKernelActivity(CUpti_ActivityKernel4 *kernel_record, ActivityTracker& tracker) {
    CHECK_KERNEL_RECORD(kernel_record);
    if (kernel_record->deviceId != tracker.prof_device_id_) {
      LOG(INFO) << "Kernel " << kernel_record->name << " @(deviceID = " << kernel_record->deviceId
                << ") is discarded. Current device is " << tracker.prof_device_id_ << ".";
      return ;
    }
    string kernel_name(kernel_record->name);
    // if (str_tolower(kernel_name).find("sync") == string::npos) {}
    // Increment the number of kernels recorded.
    tracker.kernel_counter_ ++;

    Kernel_t kernel(*kernel_record);
    uint64_t kernel_duration = kernel_record->end - kernel_record->start;

    int pos = FindKernelConfig(tracker.kernels_vec_sptr_, kernel);
    if (pos == -1) {            // For the first time of recording a kernel.
      kernel.invocations_ = 1;
      kernel.duration_ = kernel_duration;
      kernel.average_exec_time_ = kernel_duration;
      tracker.kernels_vec_sptr_->push_back(kernel);
    } else {                    // Update the information of a kernel recorded already.
      tracker.kernels_vec_sptr_->at(pos).invocations_ ++;
      auto diff = UINT64_MAX - tracker.kernels_vec_sptr_->at(pos).duration_;
      if (diff < kernel_duration) {
        LOG(FATAL) << "OVERFLOWED.";
      }
      tracker.kernels_vec_sptr_->at(pos).duration_ += kernel_duration;
      tracker.kernels_vec_sptr_->at(pos).average_exec_time_ =
        std::ceil(tracker.kernels_vec_sptr_->at(pos).duration_ /
                  static_cast<double>(tracker.kernels_vec_sptr_->at(pos).invocations_));
    }
    // LOG(INFO) << "END: " << kernel_record->end << ", START: " << kernel_record->start
    //           << ", DURATION: " << kernel_duration;

    Timestamp_t timestamp;
    timestamp.name_ = std::string(kernel_record->name);
    timestamp.start_ = kernel_record->start - tracker.start_timestamp_;
    timestamp.end_ = kernel_record->end - tracker.start_timestamp_;
    timestamp.stream_id_ = kernel_record->streamId;
    if (tracker.kernel_counter_ < tracker.timestamp_vec_sptr_->size()) {
      tracker.timestamp_vec_sptr_->at(tracker.kernel_counter_) = timestamp;
    } else {
      tracker.timestamp_vec_sptr_->push_back(timestamp);
    }
  }

  void ParseOverheadActivity(CUpti_ActivityOverhead *overhead_record, ActivityTracker& tracker) {
    if (ActivityTracker::Get().GetBuffOverheadRecord()) {
      sstream temp_ss;

      // Updated.
      // LOG(INFO) << "OVERHEAD: " << getActivityOverheadKindString(overhead_record->overheadKind) << ", "
      //           << static_cast<double>((overhead_record->end - overhead_record->start) / NS2MS) << " ms, "
      //           << getActivityObjectKindString(overhead_record->objectKind);
      temp_ss << getActivityOverheadKindString(overhead_record->overheadKind) << ","
              << ((overhead_record->end - overhead_record->start) / NS2MS)
              << " ms," << getActivityObjectKindString(overhead_record->objectKind) << std::endl;

      InfoLog::Get().RecordInfoLog("CuptiOverhead", "cupti_overhead", "CUPTI-OVERHEAD", temp_ss.str());

      temp_ss.str("");
      temp_ss.clear();
    }
  }

  void checkNullTerminator(char *str, size_t len, size_t max_len) {
    if (len >= max_len) {
      str[max_len - 1] = '\0';
    }
  }

  void ParseMetricActivity(CUpti_ActivityMetric *metric_record, ActivityTracker& tracker) {
    Metric_t curr_metric;
    size_t size = NAME_SHORT;
    size_t val_kind_size = sizeof(CUpti_MetricValueKind);

    curr_metric.metric_id_ = metric_record->id;
    CUPTI_CHECK(cuptiMetricGetAttribute(curr_metric.metric_id_,
                                        CUPTI_METRIC_ATTR_NAME,
                                        &size,
                                        reinterpret_cast<uint8_t*>(curr_metric.metric_name_)));
    checkNullTerminator(curr_metric.metric_name_, size, NAME_SHORT);

    CUPTI_CHECK(cuptiMetricGetAttribute(curr_metric.metric_id_,
                                        CUPTI_METRIC_ATTR_VALUE_KIND,
                                        &val_kind_size, &curr_metric.val_kind_));

    if (metric_record->flags == CUPTI_ACTIVITY_FLAG_METRIC_VALUE_INVALID) {
      LOG(FATAL) << "The value(s) required to calculate the metric "
                 << curr_metric.metric_name_ << " is MISSING!";
    } else if (metric_record->flags == CUPTI_ACTIVITY_FLAG_METRIC_OVERFLOWED) {
      LOG(FATAL) << "The value of metric " << curr_metric.metric_name_ << " OVERFLOWED!";
    }

    switch(curr_metric.val_kind_) {
    case CUPTI_METRIC_VALUE_KIND_DOUBLE:
      curr_metric.double_val_ = metric_record->value.metricValueDouble;
      LOG(INFO) << "Metric " << curr_metric.metric_name_ << " = " << curr_metric.double_val_;
      break;
    case CUPTI_METRIC_VALUE_KIND_UINT64:
      curr_metric.uint64_val_ = metric_record->value.metricValueUint64;
      LOG(INFO) << "Metric " << curr_metric.metric_name_ << " = " << curr_metric.uint64_val_;
      break;
    case CUPTI_METRIC_VALUE_KIND_INT64:
      curr_metric.int64_val_ = metric_record->value.metricValueInt64;
      LOG(INFO) << "Metric " << curr_metric.metric_name_ << " = " << curr_metric.int64_val_;
      break;
    case CUPTI_METRIC_VALUE_KIND_PERCENT:
      curr_metric.percent_val_ = metric_record->value.metricValuePercent;
      LOG(INFO) << "Metric " << curr_metric.metric_name_ << " = " << curr_metric.percent_val_ << "%";
      break;
    case CUPTI_METRIC_VALUE_KIND_THROUGHPUT:
      curr_metric.throughput_val_ = metric_record->value.metricValueThroughput;
      LOG(INFO) << "Metric " << curr_metric.metric_name_ << " = " << curr_metric.throughput_val_
                << " bytes/sec";
      break;
    case CUPTI_METRIC_VALUE_KIND_UTILIZATION_LEVEL:
      curr_metric.utilization_val_ = metric_record->value.metricValueUtilizationLevel;
      {
        string utilization_level;
        switch (metric_record->value.metricValueUtilizationLevel) {
        case CUPTI_METRIC_VALUE_UTILIZATION_IDLE:
          utilization_level = "IDLE";
          break;
        case CUPTI_METRIC_VALUE_UTILIZATION_LOW:
          utilization_level = "LOW";
          break;
        case CUPTI_METRIC_VALUE_UTILIZATION_MID:
          utilization_level = "MID";
          break;
        case CUPTI_METRIC_VALUE_UTILIZATION_HIGH:
          utilization_level = "HIGH";
          break;
        case CUPTI_METRIC_VALUE_UTILIZATION_MAX:
          utilization_level = "MAX";
          break;
        default:
          utilization_level = "UNKNOWN";
        }
        LOG(INFO) << "Metric " << curr_metric.metric_name_ << " = utilization level " << utilization_level;
      }
      break;
    default:
      LOG(FATAL) << "Unknown metric value kind.";
    }
  }

  void ParseCUptiActivity(CUpti_Activity *record, ActivityTracker& tracker) {
    // Only parse kernel information.
    if (record->kind == CUPTI_ACTIVITY_KIND_KERNEL or
        record->kind == CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL) { // Handle kernel profiling information.
      ParseKernelActivity(reinterpret_cast<CUpti_ActivityKernel4 *>(record), tracker);
    } else if (record->kind == CUPTI_ACTIVITY_KIND_OVERHEAD) { // Handle profiling overhead information.
      ParseOverheadActivity(reinterpret_cast<CUpti_ActivityOverhead *>(record), tracker);
    } else if (record->kind == CUPTI_ACTIVITY_KIND_METRIC) { // Handle metric information.
      ParseMetricActivity(reinterpret_cast<CUpti_ActivityMetric *>(record), tracker);
    }
  }

  ActivityTracker& ActivityTracker::Get() {
    // if (activity_tracker_ == NULL) {
    //   boost::mutex::scoped_lock instance_lock(s_prof_mutex);
    //   if (activity_tracker_ == NULL) {
    //     activity_tracker_ = new ActivityTracker();
    //   }
    // }
    static ActivityTracker activity_tracker_;

    return activity_tracker_;
  }

  ActivityTracker::ActivityTracker() {
    // Initialize the kernel_counter_ variable, which is used to record then
    // number of kernels profiled currently.
    this->kernel_counter_ = 0;
    this->config_flag_ = false;
    this->prof_device_id_ = -1;
    this->start_timestamp_ = 0;
    this->kernel_launch_overhead_ = 0;

    this->kernels_vec_sptr_ = boost::make_shared<std::vector<Kernel_t>>();
    this->timestamp_vec_sptr_ = boost::make_shared<std::vector<Timestamp_t>>();

    // Initialize the seg_tree_-related variables.
    this->next_node_idx_ = 0;
    this->tree_nodes_count_ = 0;
    this->seg_tree_ = nullptr;

    this->buff_overhead_record_ = true;
    this->init_flag_ = false;

    // Disable the profiler.
    s_profiler_flag = false;
  }

  ActivityTracker::~ActivityTracker() {
    // Empty the information vector of kernels recorded and timestamps of them.
    if (!kernels_vec_sptr_->empty()) {
      this->kernels_vec_sptr_->clear();
    }
    if (!timestamp_vec_sptr_->empty()) {
      this->timestamp_vec_sptr_->clear();
    }

    // Empty the seg_tree_ which is used to calculate the idle time of the GPU device.
    if (seg_tree_ != nullptr) {
      delete[] this->seg_tree_;
    }
    this->next_node_idx_ = 0;
    this->tree_nodes_count_ = 0;
    this->seg_tree_ = nullptr;
  }

  void ActivityTracker::Init() {
    if (init_flag_) {
      LOG(INFO) << "The activity tracker has already been initialized.";

      return ;
    }

    // In default situation, the resource tracker is disabled.
    this->prof_device_id_ = -1;
    s_profiler_flag = false;    // Disable the profiler.

    // Enable tracking concurrent kernels' information
    CUPTI_CHECK(cuptiActivityEnable(CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL));
    CUPTI_CHECK(cuptiActivityEnable(CUPTI_ACTIVITY_KIND_OVERHEAD));
    // CUPTI_CHECK(cuptiActivityEnable(CUPTI_ACTIVITY_KIND_METRIC));

    if (!this->config_flag_) {
      // Register functions for requesting buffer or processing buffer.
      CUPTI_CHECK(cuptiActivityRegisterCallbacks(BufferRequested, BufferCompleted));

      // Double the buffer size and the number of buffers.
      size_t buffer_size = 0, buffer_val_size = sizeof(size_t);
      CUPTI_CHECK(cuptiActivityGetAttribute(CUPTI_ACTIVITY_ATTR_DEVICE_BUFFER_SIZE, &buffer_val_size, &buffer_size));
      buffer_size *= 2;
      CUPTI_CHECK(cuptiActivitySetAttribute(CUPTI_ACTIVITY_ATTR_DEVICE_BUFFER_SIZE, &buffer_val_size, &buffer_size));

      CUPTI_CHECK(cuptiActivityGetAttribute(CUPTI_ACTIVITY_ATTR_DEVICE_BUFFER_POOL_LIMIT, &buffer_val_size, &buffer_size));
      buffer_size *= 2;
      CUPTI_CHECK(cuptiActivitySetAttribute(CUPTI_ACTIVITY_ATTR_DEVICE_BUFFER_POOL_LIMIT, &buffer_val_size, &buffer_size));

      this->config_flag_ = true;
    }

    init_flag_ = true;
    LOG(INFO) << "The initialization of the asynchronous resource tracker is COMPLETED!";
  }

  void ActivityTracker::Destroy() {
    if (this->config_flag_) {
      // Disable all activities.
      CUPTI_CHECK(cuptiActivityDisable(CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL));
      CUPTI_CHECK(cuptiActivityDisable(CUPTI_ACTIVITY_KIND_OVERHEAD));
      // CUPTI_CHECK(cuptiActivityDisable(CUPTI_ACTIVITY_KIND_METRIC));

      size_t buffer_size = 0, buffer_val_size = sizeof(size_t);

      CUPTI_CHECK(cuptiActivityGetAttribute(CUPTI_ACTIVITY_ATTR_DEVICE_BUFFER_SIZE, &buffer_val_size, &buffer_size));
      buffer_size /= 2;
      CUPTI_CHECK(cuptiActivitySetAttribute(CUPTI_ACTIVITY_ATTR_DEVICE_BUFFER_SIZE, &buffer_val_size, &buffer_size));

      CUPTI_CHECK(cuptiActivityGetAttribute(CUPTI_ACTIVITY_ATTR_DEVICE_BUFFER_POOL_LIMIT, &buffer_val_size, &buffer_size));
      buffer_size /= 2;
      CUPTI_CHECK(cuptiActivitySetAttribute(CUPTI_ACTIVITY_ATTR_DEVICE_BUFFER_POOL_LIMIT, &buffer_val_size, &buffer_size));

      this->config_flag_ = false;
    }

    // Cleanup CUPTI.
    cuptiFinalize();

    init_flag_ = false;
    LOG(INFO) << "ActivityTracker object has been DESTROYED!";
  }

  void ActivityTracker::Lock() {
    s_prof_mutex.lock();
  }

  void ActivityTracker::Unlock() {
    s_prof_mutex.unlock();
  }

  void ActivityTracker::Start(int device_id) {
    // Initialize the device needed to be profiled.
    this->prof_device_id_ = device_id;

    this->kernel_launch_overhead_ = 0; // Initialize the launch overhead to 0.

    if (!this->kernels_vec_sptr_->empty()) {
      this->kernels_vec_sptr_->clear();
    }
    // Initialize the vector storing kernels information.
    this->kernels_vec_sptr_ = boost::make_shared<std::vector<Kernel_t>>();

    if (!this->timestamp_vec_sptr_->empty()) {
      this->timestamp_vec_sptr_->clear();
    }
    // Initialize the vector that store the kernels' timestamp to a object-specific value.
    this->timestamp_vec_sptr_ = boost::make_shared<std::vector<Timestamp_t>>();

    // Get the start timestamp of profiling record, which is used to normalize kernel's timestamps.
    CUPTI_CHECK(cuptiGetTimestamp(&this->start_timestamp_));

    // Enable the GPU profiler after all preparations have been done.
    s_profiler_flag = true;     // Start the ActivityTracker.
    LOG(INFO) << "START profiling on the " << this->prof_device_id_ << "-th GPU device!";
  }

  void ActivityTracker::Stop() {
    CUDA_CHECK(cudaDeviceSynchronize()); // Synchronization.
    CUPTI_CHECK(cuptiActivityFlushAll(0)); // Flush the CUPTI buffer.

    // Stop the profiling.
    s_profiler_flag = false;
    LOG(INFO) << "STOP profiling on device " << this->prof_device_id_ << "! TOTAL " << this->kernel_counter_ << " kernels are recorded!";
  }

  boost::shared_ptr<std::vector<Kernel_t>> ActivityTracker::GetKernelsRecorded() {
    return this->kernels_vec_sptr_;
  }

  uint64_t ActivityTracker::GetKernelLaunchOverhead() {
    unsigned int min_invocations = std::numeric_limits<unsigned int>::max();
    for (auto kernel_rec : (*kernels_vec_sptr_)) {
      min_invocations = std::min(min_invocations, kernel_rec.invocations_);
      // LOG(INFO) << "Kernel name: " << kernel_rec.name_ << " [" << kernel_rec.invocations_ << "].";
    }
    std::sort(timestamp_vec_sptr_->begin(), timestamp_vec_sptr_->end(), Compare<Timestamp_t>);

    unsigned int kernels_per_iter = timestamp_vec_sptr_->size() / min_invocations;
    // To avoid additional kernel recorded.

    uint64_t total_launch_overhead = 0;
    if (kernels_per_iter > 1) {
      for (int i = 0; i < min_invocations; ++ i) {
        for (int j = 1; j < kernels_per_iter; ++ j) {
          total_launch_overhead += (timestamp_vec_sptr_->at(i * kernels_per_iter + j).start_ -
                                    timestamp_vec_sptr_->at(i * kernels_per_iter + j - 1).end_);
        }
      }
    } else {
      for (int i = 0; i < (min_invocations - 1); ++ i) {
        total_launch_overhead += (timestamp_vec_sptr_->at(i + 1).start_ - timestamp_vec_sptr_->at(i).end_);
      }
    }

    // LOG(INFO) << "total_launch_overhead = " << total_launch_overhead << "; kernels_per_iter = "
    //           << kernels_per_iter << "; min_invocations = " << min_invocations;
    if (kernels_per_iter > 1) {
      kernel_launch_overhead_ = total_launch_overhead / ((kernels_per_iter - 1) * min_invocations);
    } else {
      kernel_launch_overhead_ = total_launch_overhead / (min_invocations - 1);
    }

    if (this->buff_overhead_record_) {
      sstream temp_ss;
      temp_ss << "timestamp_vec_buffer," << sizeof(Timestamp_t) * timestamp_vec_sptr_->size() << ","
              << "kernel_temp_buffer," << sizeof(Kernel_t) * kernels_vec_sptr_->size() << std::endl;
      InfoLog::Get().RecordInfoLog("CuptiOverhead", "buffer_overhead", "TEMP-BUFFER-OVERHEAD", temp_ss.str());

      // LOG(INFO) << "The launch overhead of a kernel is " << kernel_launch_overhead_ << "!";
      temp_ss.str("");
      temp_ss.clear();
    }

    return kernel_launch_overhead_;
  }

  void ActivityTracker::TreeBuild(SegTreeNode_ptr seg_tree, int node_id, int start, int end) {
    int mid_1 = (start + end) >> 1;
    int mid_2 = mid_1;

    seg_tree[node_id].start_ = start;
    seg_tree[node_id].end_ = end;
    seg_tree[node_id].c_flag_ = false;
    seg_tree[node_id].left_child_ = seg_tree[node_id].right_child_ = 0;

    if (start == end) {
      return ;
    }
    if (start == mid_1) {
      mid_1 = start;
      mid_2 = end;
    }

    // Build the segment tree in a recursive way.
    seg_tree[node_id].left_child_ = this->next_node_idx_ ++;
    TreeBuild(seg_tree, seg_tree[node_id].left_child_, start, mid_1);
    seg_tree[node_id].right_child_ = this->next_node_idx_ ++;
    TreeBuild(seg_tree, seg_tree[node_id].right_child_, mid_2, end);
  }

  void ActivityTracker::TreeInsert(SegTreeNode_ptr seg_tree, int current_node, int start, int end) {
    if ((start >= end) or (seg_tree[current_node].c_flag_)) {
      return ;
    }

    int start_val = seg_tree[current_node].start_;
    int end_val = seg_tree[current_node].end_;

    // Termination condition.
    if (start_val == start and end_val == end) {
      seg_tree[current_node].c_flag_ = true;
      return ;
    }

    int mid_val = (start_val + end_val) >> 1;
    if (end <= mid_val) {
      TreeInsert(seg_tree, seg_tree[current_node].left_child_, start, end);
    } else if (start >= mid_val) {
      TreeInsert(seg_tree, seg_tree[current_node].right_child_, start, end);
    } else {
      TreeInsert(seg_tree, seg_tree[current_node].left_child_, start, mid_val);
      TreeInsert(seg_tree, seg_tree[current_node].right_child_, mid_val, end);
    }

    return ;
  }

  uint64_t ActivityTracker::TreeTraverse(const SegTreeNode_ptr seg_tree, const vector<uint64_t>& time_vec, int node_id) {
    if (seg_tree[node_id].c_flag_) {
      return (time_vec[seg_tree[node_id].end_] - time_vec[seg_tree[node_id].start_]);
    }

    uint64_t left_value = 0, right_value = 0;
    if (seg_tree[node_id].left_child_ != 0) {
      left_value = TreeTraverse(seg_tree, time_vec, seg_tree[node_id].left_child_);
    }
    if (seg_tree[node_id].right_child_ != 0) {
      right_value = TreeTraverse(seg_tree, time_vec, seg_tree[node_id].right_child_);
    }

    return (left_value + right_value);
  }

  void ActivityTracker::ComputeOccupancyRatio(const string layer_name, const unsigned int parallel_degree) {
    if (this->kernel_counter_ == 0) {
      LOG(INFO) << "There is no kernel recorded.";
      return ;
    }

    vector<uint64_t> time_vec;
    // Remove redundant timestamp from original record.
    for (int i = 0; i < this->kernel_counter_; ++ i) {
      auto start_iter = std::find(time_vec.begin(), time_vec.end(), this->timestamp_vec_sptr_->at(i).start_);
      if (start_iter == time_vec.end()) { // Only push new time point to the time_vec.
        time_vec.push_back(timestamp_vec_sptr_->at(i).start_);
      }
      auto end_iter = std::find(time_vec.begin(), time_vec.end(), this->timestamp_vec_sptr_->at(i).end_);
      if (end_iter == time_vec.end()) {
        time_vec.push_back(timestamp_vec_sptr_->at(i).end_);
      }
    }

    std::sort(time_vec.begin(), time_vec.end(), Compare<uint64_t>);
    // Allocate memory for seg_tree.
    if (this->tree_nodes_count_ < (this->kernel_counter_ * 2 * 4)) {
      delete[] this->seg_tree_;
      this->tree_nodes_count_ = this->kernel_counter_ * 2 * 4;
      this->seg_tree_ = new SegTreeNode_t[this->tree_nodes_count_];

      if (this->seg_tree_ == nullptr) {
        LOG(FATAL) << "Failed to allocate memory for the segment tree!";
      } else {
        // LOG(INFO) << "ALLOCATE " << this->tree_nodes_count_ << " nodes!";
      }
    }
    std::memset(this->seg_tree_, 0, this->tree_nodes_count_ * sizeof(SegTreeNode_t));

    // Segment tree construction.
    this->next_node_idx_ = 1;   // Index that tree nodes constructed should be stored from.
    TreeBuild(this->seg_tree_, 0, 0, time_vec.size() - 1);

    // Fill all busy time slots.
    for (int i = 0; i < this->kernel_counter_; ++ i) {
      int start_pos = std::find(time_vec.begin(), time_vec.end(),
                                this->timestamp_vec_sptr_->at(i).start_) - time_vec.begin();
      int end_pos = std::find(time_vec.begin(), time_vec.end(),
                              this->timestamp_vec_sptr_->at(i).end_) - time_vec.begin();
      TreeInsert(this->seg_tree_, 0, start_pos, end_pos);
    }

    // Get the total busy time in a particular time range.
    uint64_t busy_time = TreeTraverse(this->seg_tree_, time_vec, 0);
    uint64_t total_time = time_vec.back() - time_vec.front();
    uint64_t idle_time = total_time - busy_time;
    double occupancy_ratio = static_cast<double>(busy_time) / static_cast<double>(total_time);
    // LOG(INFO) << "Occupancy ratio of layer " << layer_name << " (#streams = " << parallel_degree
    //           << ") is " << occupancy_ratio << ". IDLE = " << idle_time << ", TOTAL = " << total_time;

    sstream filename_ss, content_ss;
    filename_ss << layer_name << "_ParallelDeg_" << parallel_degree << "_occ";
    content_ss << layer_name << "," << parallel_degree << "," << occupancy_ratio << "," << total_time
               << "," << idle_time << std::endl;
    InfoLog::Get().RecordInfoLog("OccupancyRatio", "Meanlingless_Value", filename_ss.str(), content_ss.str());
    filename_ss.str("");
    content_ss.str("");
    filename_ss.clear();
    content_ss.clear();

    // Free memory allocated.
    time_vec.clear();

    return ;
  }

  void ActivityTracker::TimestampLog(const string filename) const {
    sstream temp_ss;

    const int step = 40;
    for (int i = 0; i < this->timestamp_vec_sptr_->size(); i += step) {
      temp_ss.str("");
      temp_ss.clear();
      for (int j = 0; j < step and (j + i) < this->timestamp_vec_sptr_->size(); ++ j) {
        const Timestamp_t *temp_timestamp = &this->timestamp_vec_sptr_->at(i + j);
        temp_ss << temp_timestamp->name_ << "," << temp_timestamp->stream_id_ << ","
                << temp_timestamp->start_ << "," << temp_timestamp->end_ << std::endl;
      }

      InfoLog::Get().RecordInfoLog("KernelInfo", "Meaningless_Value", filename + "_timestamp", temp_ss.str());
    }

    temp_ss.str("");
    temp_ss.clear();
  }

  void ActivityTracker::TempBufRelease() {
    if (!this->kernels_vec_sptr_->empty()) {
      this->kernels_vec_sptr_->clear();
    }
    this->kernel_counter_ = 0;
    if (!this->timestamp_vec_sptr_->empty()) {
      this->timestamp_vec_sptr_->clear();
    }
    if (this->seg_tree_ != nullptr) {
      delete[] this->seg_tree_;
      this->seg_tree_ = nullptr;
      this->tree_nodes_count_ = 0;
    }
  }
} /* namespace caffe */

#endif // USE_PROF
#endif /* CPU_ONLY */
